
#include "../devmap.h"
#include "../queue.h"

static const char dummy_code[] =
{
	"__kernel void dummy( __global Real *a, Real x )\n"
	"{ a[ get_global_id(0) ] = x*get_global_id(0); }"
};


int main(int argc, char *argv[])
{
	int         result     = 0;
	ocl_driver *OpenCL     = NULL;
	ocl_devmap *devmap     = NULL;
	ocl_context context    = OCL_CONTEXT_INIT;
	ocl_code    sources    = OCL_CODE_INIT;
	ocl_program program    = OCL_PROGRAM_INIT;
	ocl_kernel  k_dummy    = OCL_KERNEL_INIT;
	ocl_kernel  k_add      = OCL_KERNEL_INIT;
	ocl_buffer  buf_a      = OCL_BUFFER_INIT;
	ocl_buffer  buf_b      = OCL_BUFFER_INIT;
	ocl_buffer  buf_c      = OCL_BUFFER_INIT;
	ocl_queue   Q          = OCL_QUEUE_INIT;
	Real       *a = NULL, *b=NULL, *c=NULL;
	size_t      num   = 100;
	size_t      wgs   = 10;
	cl_uint     i, j;
	
	/* I need a device map ! */
	if( argc <= 1 )
	{
		fprintf( stderr, "usage: devmap [-v]\n");
		return -1;
	}
	
	/* I may be verbose */
	if( argc > 2 && strcmp( argv[2], "-v" ) == 0 )
		ocl_logfile = stderr;
	
	/* load OpenCL infos */
	OpenCL = ocl_driver_load();
	if( ! OpenCL )
	{
		ocl_perror( stderr, "ocl_driver_init");
		return -1;
	}
	
	/* parse the devmap */
	devmap = ocl_devmap_create( OpenCL, argv[1] );
	if( !devmap )
	{
		mds_perror( stderr, "ocl_devmap_create(%s)", argv[1]);
		result = -1;
		goto END;
	}
	
	ocl_devmap_write(devmap, OpenCL, stderr );
	
	/* create a context */
	if( ocl_context_map( &context, devmap) < 0 )
	{
		ocl_perror( stderr, "ocl_context_map");
		result = -1;
		goto END;
	}
	
	/* load some sources */
	if( ocl_code_read( &sources, "sample.cl" ) < 0 )
	{
		ocl_perror( stderr, "ocl_code_read" );
		result = -1;
		goto END;
	}
	
	if( ocl_code_copy( &sources, dummy_code, strlen( dummy_code ) ) <  0 )
	{
		ocl_perror( stderr, "ocl_code_copy" );
		result = -1;
		goto END;
	}
	
	
	
	/* create a program */
	if( ocl_program_load_sources( &program, &context, &sources ) < 0 )
	{
		ocl_perror( stderr, "ocl_program_load_sources" );
		result = -1;
		goto END;
	}
	
	/* compile it, with predefined options */
	if( ocl_program_compile( &program, ocl_options ) < 0 )
	{
		ocl_perror( stderr, "ocl_program_compilation failure !");
		ocl_program_display_build_logs( &program, stderr );
		result = -1;
		goto END;
	}
	else {
		ocl_program_display_build_logs( &program, stderr );
	}
	
	
	
	
	/* get kernels */
	if( ocl_kernel_load( &k_dummy, &program, "dummy" ) < 0 )
	{
		ocl_perror( stderr, "ocl_kernel_load failure !");
		result = -1;
		goto END;
	}
	
	if( ocl_kernel_load( &k_add, &program, "add" ) < 0 )
	{
		ocl_perror( stderr, "ocl_kernel_load failure !");
		result = -1;
		goto END;
	}
	
	/* create buffers */
	if(ocl_buffer_of(Real, &buf_a, &context, CL_MEM_READ_WRITE, num, NULL ) < 0  ||
	   ocl_buffer_of(Real, &buf_b, &context, CL_MEM_READ_WRITE, num, NULL ) < 0  || 
	   ocl_buffer_of(Real, &buf_c, &context, CL_MEM_READ_WRITE, num, NULL ) < 0 )
	{
		ocl_perror( stderr, "creating buffer");
		result = -1;
		goto END;
	}
	
	/* assign arguments */
	OCL_KERNEL_SET_ARGS_INIT(&k_dummy,arg_index,res);
	
	const Real x = 1.0;
	OCL_KERNEL_SET_ARG_BUFFER( &buf_a );
	OCL_KERNEL_SET_ARG(Real,x);
	
	if( res < 0 ) 
	{
		ocl_perror( stderr, "Failure for %s.arg[%u]", k_dummy.FUNCTION_NAME, arg_index);
		result = -1;
		goto END;
	}
	
	OCL_KERNEL_SET_ARGS_DONE();
	
	OCL_KERNEL_SET_ARGS_INIT(&k_add,arg_index,res);
	OCL_KERNEL_SET_ARG_BUFFER( &buf_a );
	OCL_KERNEL_SET_ARG_BUFFER( &buf_b );
	OCL_KERNEL_SET_ARG_BUFFER( &buf_c );
	if( res < 0 ) 
	{
		ocl_perror( stderr, "Failure for %s.arg[%u]", k_add.FUNCTION_NAME, arg_index);
		result = -1;
		goto END;
	}
	OCL_KERNEL_SET_ARGS_DONE();
	
	/* create user space buffers */
	a = MDS_CALLOC_AS(Real,num);
	b = MDS_CALLOC_AS(Real,num);
	c = MDS_CALLOC_AS(Real,num);
	if( !a || !b || !c )
	{
		mds_perror(stderr, "memory failure");
		result = -1;
		goto END;
	}
	
	for(  i=0; i < context.num_devices; ++i )
	{
		const cl_device_id d = context.DEVICES[i];
		const ocl_device  *D = ocl_driver_device( d );
		fprintf( stderr, "running on '%s'\n", D->NAME ); 
		if( ocl_queue_load( &Q, &context, d, 0 ) < 0 )
		{
			ocl_perror( stderr, "couldn't create queue");
			continue;
		}
		const size_t global_work_size[4] = { num };
		size_t       local_work_size[4]  = { wgs };
		
		fprintf( stderr, "Executing '%s'\n", k_dummy.FUNCTION_NAME );
		if( ocl_enqueue_ndrange_kernel( &Q, &k_dummy, 1, NULL, global_work_size, local_work_size, OCL_NO_EVENTS ) < 0 )
		{
			ocl_perror( stderr, "ndrange '%s'", k_dummy.FUNCTION_NAME );
			goto DONE;
		}
		
		fprintf( stderr, "Reading a\n");
		if( ocl_enqueue_read_buffer( &Q, &buf_a, CL_TRUE, 0, buf_a.SIZE, a, OCL_NO_EVENTS ) < 0 )
		{
			ocl_perror( stderr, "reading a");
			goto DONE;
		}
		
		for( j=0; j < num; ++j )
		{
			if( (int)a[j] != (int)j )
			{
				fprintf(stderr, "a reading: bad value !!\n");
				goto DONE;
			}
			b[j] = num-j;
		}
		
		fprintf( stderr, "Writing b\n");
		if( ocl_enqueue_write_buffer( &Q, &buf_b, CL_TRUE, 0, buf_b.SIZE, b, OCL_NO_EVENTS ) < 0 )
		{
			ocl_perror( stderr, "writing b");
			goto DONE;
		}
		
		fprintf( stderr, "Executing '%s'\n", k_add.FUNCTION_NAME );
		if( ocl_enqueue_ndrange_kernel( &Q, &k_add, 1, NULL, global_work_size, local_work_size, OCL_NO_EVENTS ) < 0 )
		{
			ocl_perror( stderr, "ndrange '%s'", k_add.FUNCTION_NAME );
			goto DONE;
		}
		
		fprintf( stderr, "Reading c\n");
		if( ocl_enqueue_read_buffer( &Q, &buf_c, CL_TRUE, 0, buf_c.SIZE, c, OCL_NO_EVENTS ) < 0 )
		{
			ocl_perror( stderr, "reading c");
			goto DONE;
		}
		
		for( j=0; j < num; ++j )
		{
			if( (int) c[j] != (int) num )
			{
				fprintf(stderr, "c reading: bad value !!\n");
				goto DONE;
			}
		}
		
	DONE:
		ocl_queue_quit( &Q );
	}
	
	
END:
	MDS_MEMORY();
	MDS_FREE_AS(Real,a,num);
	MDS_FREE_AS(Real,b,num);
	MDS_FREE_AS(Real,c,num);
	
	ocl_buffer_quit( &buf_c );
	ocl_buffer_quit( &buf_b );
	ocl_buffer_quit( &buf_a );
	ocl_kernel_quit( &k_add );
	ocl_kernel_quit( &k_dummy );
	ocl_devmap_delete( devmap  );
	
	ocl_program_quit( &program );
	
	
	ocl_code_quit(    &sources  );
	
	ocl_context_quit( &context );
	ocl_driver_quit();
	MDS_MEMORY();
	
	return result;
	
	
	
}
